#include "SMO_Include.cl"
#pragma OPENCL EXTENSION cl_khr_fp64 : enable

__attribute__((reqd_work_group_size(thread_group_size, 1, 1)))
__kernel void SVMOutput(__global svm_precision* b_alpha, __global svm_precision* b_class, __global svm_precision* b_outputBuffer, 
						__global svm_precision* b_inputBuffer, __global unsigned int* b_trainingIndices, 
						__global unsigned int* b_kernelEvalBuffer, __global svm_precision* b_selfProd, 
						__constant SharedBuffer* constants){
	__local evalStruct sdata[thread_group_size];
    // load shared mem
    unsigned int tid = get_global_id(0);
	unsigned int threadID = get_local_id(0);
	if(tid < constants->cb_instanceCount){
		unsigned int n = constants->cb_instanceLength;
		unsigned int outputInd = b_kernelEvalBuffer[tid];

		sdata[threadID].ind1 = b_trainingIndices[constants->cb_ind1] * n;
		sdata[threadID].ind2 = b_trainingIndices[outputInd] * n;
	
		svm_precision result1 = 0;
		unsigned int i = 0;
		do{
			if(i != constants->cb_classIndex)
				result1 += b_inputBuffer[i+sdata[threadID].ind1] * b_inputBuffer[i+sdata[threadID].ind2];
			i++;
		}while (i < n);

		b_outputBuffer[tid] = evaluateKernel(result1,sdata[threadID].ind1/n,sdata[threadID].ind2/n,b_selfProd,constants);
		//b_outputBuffer[tid] = evaluateKernel(result1,sdata[threadID].ind1/n,sdata[threadID].ind2/n,b_selfProd,constants) * b_class[outputInd] * b_alpha[outputInd];
	}
}